{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "_lvvyWhmZkgW"
   },
   "source": [
    "# CUDA C/C++에서 가속 애플리케이션을 위한 비동기 스트리밍과 비주얼 프로파일링\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "T8l_LuQaE3JA"
   },
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "b9yWEwoOFC4B"
   },
   "source": [
    "CUDA 툴킷에는 가속화 CUDA 애플리케이션 개발을 지원하는 강력한 GUI 애플리케이션인 **NVIDIA 비주얼 프로파일러**(또는 **nvvp**)가 포함되어 있습니다. nvvp는 가속화 애플리케이션에서 일어나는 CUDA API 호츨, 커널 실행, 메모리 동작, **CUDA 스트림** 사용 상황등에 대한 정보를 타임라인을 통해 그래픽으로 보여줍니다. \n",
    "\n",
    "게다가 nvvp에는 가속화 에플리케이션을 어떻게 최적화하는 것이 가장 좋을지에 대해 추천을 하는 분석 도구도 함께 제공됩니다. nvvp에 익숙해지는 일은 CUDA 개발자에게는 필수적이라고 할 수 있습니다.\n",
    "\n",
    "본 강좌에서 여러분은 nvvp 타임라인의 도움을 받아 가면서 가속화 에플리케이션 최적화를 수행해 나갈 것입니다. 그리고 여러분의 작업을 도울 중급 CUDA 프로그래밍 기법을 베울 것입니다. 그 기법들은 **수동(non-managed or manual) 메모리 할당과 마이그레이션**, 호스트 메모리 **고정(pinning)**(또는 호스트 메모리 **페이지 잠금(page-locking)**), **비(非)디폴트 동시성 CUDA 스트림** 등입니다.\n",
    "\n",
    "본 강좌가 끝나면 여러분에게 n-body 시뮬레이터의 가속화와 최적화 과제가 주어질텐데, 이를 통해 본 강좌에서 여러분이 습득한 기법을 펼쳐 보일 수 있을 것입니다. 시뮬레이터의 동작을 올바르게 유지하면서 가속화에 성공한다면 여러분의 역량을 증명하기 위한 인증서를 받으실 수 있습니다.\n",
    "\n",
    "---\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "k8-EqoOMFC64"
   },
   "source": [
    "## 선행지식\n",
    "\n",
    "강좌의 이해를 위하여 여러분은 다음 사항을 알고 있어야 합니다.\n",
    "\n",
    "* CPU 함수 호출과 GPU 커널 구동을 포함한 C/C++ 프로그램의 작성, 컴파일, 실행하기\n",
    "* 실행 설정을 이용한 병렬 스레드 계층구조 제어하기\n",
    "* 직렬 반복문을 GPU에서 병렬로 실행하도록 코드 수정하기\n",
    "* 통합 메모리의 할당과 해제\n",
    "* 페이지 폴트와 데이터 마이그레이션에 관련된 통합 메모리의 작동 이해\n",
    "* 페이지 폴트와 데이터 마이그레이션을 줄이기 위한 비동기 메모리 프리패칭 사용\n",
    "\n",
    "---"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "150x6QRzFC9m"
   },
   "source": [
    "\n",
    "\n",
    "## 학습목표\n",
    "\n",
    "강좌를 마치면 여러분은 다음 사항을 수행할 수 있게 될 것입니다.\n",
    "\n",
    "* **NVIDIA 비주얼 프로파일러(nvvp)**를 이용하여 가속화 애플리케이션의 타임라인을 시각적으로 프로파일링하기\n",
    "* nvpp를 사용하여 가속화 에플리케이션의 최적화 기회를 인식하고 활용하기\n",
    "* CUDA 스트림을 이용하여 가속화 애플리케이션에서 커널들을 동시에 실행하기\n",
    "* **(옵션으로 주어진 고급 주제)** 고정(pinned) 메모리를 포함한 수동 메모리 할당을 이용하여 동시성 CUDA 스트림에서 비동기로 데이터 옮기기"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "1buatEkGE3uM"
   },
   "source": [
    "\n",
    "\n",
    "---\n",
    "\n",
    "## NVIDIA 비주얼 프로파일러 세팅하기\n",
    "\n",
    "[nvvp 실행 링크](/novnc)를 클릭하여 nvpp 세션을 다른 탭에서 여세요. 패스워드를 `cuda`로 입력하고 연결하면 nvpp에 접속됩니다. 다음 섹션에서 이를 이용하여 CUDA 코드를 프로파일링할 것입니다.\n",
    "\n",
    "**주의: 터치스크린 랩탑에서 윈도우즈를 사용하신다면 사용에 문제가 있을 수 있습니다. 그러할 경우에는 파이어폭스 브라우저를 사용하세요.**\n",
    "\n",
    "워크스페이스를 사용하라는 요청이 있을 경우, 그냥 기본(default) 워크스페이스를 선택하세요. nvvp가 곧 열릴 것입니다.\n",
    "\n",
    "사용하시다가 nvpp가 타임아웃이 걸릴 경우, 연결(connect) 버튼을 누르시면 다시 연결됩니다."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "E7Pny3AkE3xC"
   },
   "source": [
    "\n",
    "\n",
    "---\n",
    "\n",
    "## nvvp를 이용하여 반복적으로 코드 수정 비교하기\n",
    "\n",
    "다음 일련의 실습들을 통해 여러분은 nvvp 타임라인 작업에 숙달될 것입니다. 여러분이 이미 잘 알고 있는 기법들을 사용하여 반복적 개선을 수행하고 프로파일링을 해 볼 것입니다. 이를 통해 애플리케이션 성능에 영향을 미치는 다양한 CUDA 프로그래밍 기법에 대한 이해를 높이실 수 있을 것입니다.\n",
    "\n",
    "본 섹션의 실습을 마치고 나면, 동시성 CUDA 스트림, 수동 CUDA 메모리 할당, 복사 등을 포함한 새로운 프로그래밍 기법의 유용성을 nvvp 타임라인을 사용하면서 익히게 될 것입니다."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "5vC17mcrE3zk"
   },
   "source": [
    "\n",
    "\n",
    "\n",
    "### 실습: 컴파일된 CUDA 코드의 타임라인 검사하기\n",
    "\n",
    "[`01-vector-add.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/01-vector-add/01-vector-add.cu)(<---- 소스 파일을 새로운 탭에서 열고 수정하려면 클릭하세요.)는 실제로 동작하는 가속화 벡터합 애플리케이션입니다. 아래 코드 실행 셀을 이용하여(CTRL + 클릭으로 실행됩니다.) 컴파일하고 실행하세요. 성공을 나타내는 메시지를 보실 것입니다."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "42UFIX1aZkgX"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch sm_70 -o vector-add-no-prefetch 01-vector-add/01-vector-add.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "dGzvlPY3Zkgb"
   },
   "source": [
    "성공적으로 컴파일한 후, 컴파일된 실행 파일을 [nvvp](/novnc)로 여시고 타임라인 윈도우를 최대화한 후, 아래 내용을 실행하세요.\n",
    "\n",
    "* `addVectorsInto` 커널의 실행 시간을 보여주는 타임라인 룰러(timeline ruler)를 만드세요.\n",
    "* 타임라인에서 CPU 페이지 폴트가 일어난 지점을 찾으세요. [소스 코드](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/01-vector-add/01-vector-add.cu)의 어느 부분에서 폴트가 일어났는지 찾으세요.\n",
    "* 디바이스에서 *데이터 마이그레이션(DtoH)* 이벤트가 일어난 지점을 타임라인에서 찾으세요. 커널 수행후 CPU 페이지 폴트가 일어난 근처일 것입니다. 왜 커널 실행전 CPU 페이지 폴트가 일어난 때가 아니라 커널 수행후 CPU 페이지 폴트 부분에서 일어났을까요?\n",
    "* 타임라인상의 GPU 페이지 폴트, HtoD 데이터 마이그레이션 이벤트, `addVectorsInto` 커널 실행들 사이에 어떤 관계가 있을까요? 왜 이들이 그런 식으로 발생했는지 소스 코드를 가지고 설명할 수 있나요?\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "8zQT9hzMFYYL"
   },
   "source": [
    "### 실습: 비동기적 프리패칭 추가하기\n",
    "\n",
    "[01-vector-add-prefetch-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/01-vector-add/solutions/01-vector-add-prefetch-solution.cu)는 앞에서 본 벡터합 애플리케이션을 다음과 같이 수정한 코드입니다. 즉, `addVectorsInto`가 사용하는 세 개의 벡터를 커널 구동 전에 [cudaMemPrefetchAsync](http://docs.nvidia.com/cuda/cuda-runtime- api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1ge8dc9199943d421bc8bc7f473df12e42)를 이용하여 비동기적으로 프리패치하여 GPU에게 전달하도록 되어 있습니다. 소스 코드를 열어서 이렇게 수정된 부분을 찾아 보세요. \n",
    "\n",
    "수정 내용을 살펴보신 후, 애플리케이션을 아래 셀을 이용하여 실행하세요. 성공했다는 메시지를 보시게 될 것입니다."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "-uqUKNzzHuSU"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch sm_70 -o vector-add-prefetch 01-vector-add/solutions/01-vector-add-prefetch-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "1wTFkusAZkgg"
   },
   "source": [
    "### 실습: 프리패칭 있는 경우와 없는 경우의 타임라인 비교하기\n",
    "\n",
    "[nvvp](/novnc)을 이용하여 컴파일된 실행 파일을 열고, 프리패칭을 추가하기 이전의 벡터합 애플리케이션 세션이 열려 있는 상태대로 두십시오. 타임라인 창을 최대화하고 아래를 수행하세요. \n",
    "\n",
    "* `addVectorsInto` 커널의 실행 시간을 보여 주는 룰러를 만드세요. 비동기적 프리패칭을 추가하기 이전의 `addVectorsInto` 커널 실행 시간과 비교할 때 어떤 변화가 있나요? \n",
    "* 타임라인의 *런타임 API* 섹션에서 `cudaMemPrefetchAsync`를 찾으세요.\n",
    "* 코드 수정 전의 타임라인의 *UM 섹션*을 보시면 커널을 실행하면서 UM이 요청되는 순간 발생한 몇 가지 *GPU 페이지 폴트*가 그룹으로 묶인 것을 볼 수 있을 것입니다. 이제 프리패칭을 구현했는데 이 페이지 폴트가 여전히 존재하나요?\n",
    "* GPU 페이지 폴트가 사라지긴 했지만, 데이터는 여전히 호스트에서 디바이스로 전달되어야 합니다. 타임라인의 *데이터 마이그레이션 (HtoD)* 섹션을 보시고 두 애플리케이션에서 마이그레이션이 어떻게 다른지 비교하세요. 얼마나 많이 발생했고, 마이그레이션에 소요된 시간은 얼마이고, `addVectorsInto` 커널 실행중 상대적으로 언제 발생했는지 비교하세요.\n",
    "* 두 애플리케이션의 전체 실행시간은 어떻게 다른가요?\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "drS3IAe1FcJ4"
   },
   "source": [
    "\n",
    "\n",
    "### 실습: 초기화가 커널에서 일어나도록 수정된 코드 프로파일링하기\n",
    "\n",
    "이전의 벡터합 애플리케이션 개발 단계에서는 벡터 데이터가 CPU에서 초기화되기 때문에 `addVectorsInto` 커널이 벡터를 다루기 전에 GPU로의 마이그레이션이 필요했습니다.\n",
    "\n",
    "이번 개발 단계의 [01-init-kernel-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/02-init-kernel/solutions/01-init-kernel-solution.cu) 코드에서는 초기화가 GPU 상에서 병렬로 이루어지도록 바뀌었습니다.\n",
    "\n",
    "이제 초기화가 GPU에서 일어나므로, 프리패칭도 벡터합 연산 이전이 아닌 초기화 이전에 일어나야 합니다. 소스 코드에서 이러한 변경이 있는 부분을 찾아 보세요.\n",
    "\n",
    "찾아 보셨으면 컴파일하고 실행하세요. 성공 메시지가 출력될 것입니다."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "jz1gvhD_Zkgg"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o init-kernel 02-init-kernel/solutions/01-init-kernel-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "U07oqeJeZkgj"
   },
   "source": [
    "컴파일된 실행 파일을 nvpp의 새로운 세션에서 열고 아래를 수행하세요.\n",
    "\n",
    "* 다음 사항에 대한 타임라인 룰러를 만드세요: 전체 실행 시간, `addVectorsInto` 커널 실행 시간, 초기화 커널 실행 시간. 애플리케이션 실행 시간과 `addVectorsInto` 커널의 실행 시간을 이전 버전 애플리케이션의 값들과 비교하세요. 어떻게 다른가요?\n",
    "* *계산(compute)* 섹션을 보세요. `addVectorsInto` 커널과 초기화 커널중에 누가 GPU 시간을 주로 소비하나요?\n",
    "* 여러분의 애플리케이션에서 다음 이벤트중 무엇이 일어났나요?\n",
    "    * CPU 페이지 폴트\n",
    "    * GPU 페이지 폴트\n",
    "    * 데이터 마이그레이션(HtoD)\n",
    "    * 데이터 마이그레이션(DtoH)\n",
    "    \n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "eJMSOxvnFelc"
   },
   "source": [
    "### 실습: 호스트로의 비동기적 프리패치를 추가한  코드 프로파일링하기\n",
    "\n",
    "커널에서 수행된 벡터합의 성공 여부는 호스트에서 검사합니다. [01-prefetch-check-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/04-prefetch-check/solutions/01-prefetch-check-solution.cu)는 검사를 위해 데이터를 다시 호스트쪽으로 비동기 프리패치하도록 코드가 수정되어 있습니다.\n",
    "\n",
    "코드 수정을 검토한 후, 컴파일하고 실행하세요. 성공 메시지가 출력될 것입니다."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "35NQN-kbZkgk"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o prefetch-to-host 04-prefetch-check/solutions/01-prefetch-check-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "P1Ro6O0yZkgn"
   },
   "source": [
    "컴파일된 실행 파일을 nvpp의 새로운 세션에서 열고 아래를 수행하세요.\n",
    "\n",
    "* 타임라인의 *통합 메모리* 섹션을 사용해서 CPU로의 프리패칭을 추가하기 이전과 이후의 *데이터 마이그레이션(DtoH)* 이벤트를 비교하세요.\n",
    "    * CPU 페이지 폴트가 어떻게 다른가요?\n",
    "    * DtoH 마이그레이션에 소요된 시간이 전반적으로 어떻게 다른가요?\n",
    "* 다음 섹션과의 부드러운 연결을 위해 타임라인의 *스트림* 섹션을 보세요. 모든 커널 수행이 *디폴트* 스트림에서 일어난 것과 그 실행이 순차적임을 확인하세요. 스트림은 다음 섹션에서 배울 것입니다.\n",
    "\n",
    "---\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "cZySaTLNFo7V"
   },
   "source": [
    "\n",
    "## 동시성 CUDA 스트림\n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 979,
     "status": "ok",
     "timestamp": 1541022345614,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "x_qBB3UwZkgo",
    "outputId": "1f20333a-90dc-40a3-87df-f71b28d36ac8"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vR-tYNce--gUNXExn_KU29u9EldZY4Jo2R_f9fccxulPH_Y5S_bUrG5Y8qyxsUuuyC8cRs9gQbo9ZTu/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "BxCtcj1JZkgq"
   },
   "source": [
    "CUDA 프로그래밍에서 **스트림(stream)**이란 순서대로 실행되는 일련의 명령들을 말합니다. CUDA 애플리케이션에서 커널 실행과 일부 메모리 전달은 CUDA 스트림 안에서 일어납니다. 지금까지 여러분은 명시적으로 CUDA 스트림을 다루지는 않았지만, 직전의 실습에서 nvvp 타임라인을 볼 때 CUDA 코드가 커널을 *디폴트 스트림*이라고 부르는 스트림 안에서 실행하는 것을 본 적이 있습니다. \n",
    "\n",
    "CUDA 프로그래머는 비(非)디폴트 CUDA 스트림을 직접 만들고 사용할 수 있습니다. 이를 통해 여러 개의 커널을 서로 다른 스트림에서 동시에 실행하는 다중 연산이 가능합니다. 여러 스트림을 사용함으로써 여러분은 가속화 애플리케이션을 병렬화하는 또 하나의 계층을 추가하게 되고, 애플리케이션 최적화의 기회는 더 늘어나게 됩니다.\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "oAzZRU0RFt6a"
   },
   "source": [
    "\n",
    "### CUDA 스트림의 동작을 지배하는 규칙\n",
    "\n",
    "스트림을 효과적으로 사용하기 위해 알아두어야 할 몇 가지 규칙이 있습니다.\n",
    "\n",
    "* 한 스트림 안에서의 실행은 순차적으로 일어납니다.\n",
    "* 서로 다른 비디폴트 스트림에서의 동작은 실행 순서가 보장되지 않습니다.\n",
    "* 디폴트 스트림은 블로킹(blocking) 성질을 가집니다. 즉, 시작하기 전에 다른 모든 스트림의 실행이 완료되기를 기다린 후 실행되며, 일단 실행되면 완료되기 전까지는 다른 스트림을 블록합니다. \n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "Nl-2CG5vFt9m"
   },
   "source": [
    "### 비디폴트 CUDA 스트림의 생성/사용/제거\n",
    "\n",
    "다음 코드 스니펫은 비디폴트 CUDA 스트림을 생성/사용/제거하는 예제입니다. CUDA 커널을 비디폴트 스트림에서 구동하기 위해서 실행 설정의 네 번째 인수로 스트림을 전달하는 것을 볼 수 있습니다. 지금까지는 두 개의 인수만을 사용해왔습니다.\n",
    "\n",
    "```cpp\n",
    "cudaStream_t stream;       // CUDA streams are of type `cudaStream_t`.\n",
    "cudaStreamCreate(&stream); // Note that a pointer must be passed to `cudaCreateStream`.\n",
    "\n",
    "someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>(); // `stream` is passed as 4th EC argument.\n",
    "\n",
    "cudaStreamDestroy(stream); // Note that a value, not a pointer, is passed to `cudaDestroyStream`.\n",
    "```\n",
    "\n",
    "본 강좌의 범위 밖이긴 하지만 언급할 만한 내용은 옵션인 세 번째 인수입니다. 이 인수는 블록당 동적으로 할당되는 **공유 메모리**(지금은 다루지 않는 상급 주제입니다.)의 바이트 수를 지정하는 데에 쓰입니다. 블록당 할당되는 공유 메모리의 디폴트값은 0이며 본 강좌에서는 네 번째 인수를 설정할 때 세 번째 인수를 항상 0으로 설정할 것입니다. 실제 관심 있는 부분은 네 번째 인수이니까요.\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "cu-cNwWCFuAm"
   },
   "source": [
    "### 실습: 디폴트 스트림 동작 예측\n",
    "\n",
    "[`01-print-numbers`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/05-stream-intro/01-print-numbers.cu) 애플리케이션은 정수값을 입력받아 출력하는 간단한 `printNumber` 커널을 가지고 있습니다. 이 커널은 하나의 스레드를 가지는 하나의 블록에서 실행되지만 다섯 번 실행되는데, for 반복문을 이용해서 구동될 때마다 몇 번째 반복인가를 인수로 전달합니다. \n",
    "\n",
    "아래 셀을 이용하여 [01-print-numbers](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/05-stream-intro/01-print-numbers.cu)을 컴파일하고 실행하세요. 숫자가 0부터 4까지 출력되는 것을 볼 수 있을 것입니다."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "zl8UyAc0Zkgs"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o print-numbers 05-stream-intro/01-print-numbers.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "b-l7E8H8Zkgv"
   },
   "source": [
    "기본적으로 커널이 디폴트 스트림에서 실행된다고 했었는데, 다섯 번의 커널 구동이 순차적으로 이루어질까요, 아니면 동시에 이루어질까요? 여러분은 답안의 근거를 밝히기 위해 두 가지 점을 언급해야 합니다. 실행 파일을 nvvp의 신규 세션에서 열고, 타임라인을 최대화하고 커널 구동 부분을 확대해서 여러분의 답을 검증하세요.\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "22vJ1TezGEbB"
   },
   "source": [
    "### 실습: 동시성 CUDA 스트림 구현하기\n",
    "\n",
    "다섯 개의 커널 모두 같은 스트림에서 실행되므로 이들이 순차적으로 실행된다는 점은 그다지 새삼스러운 일이 아닐 것입니다. 또한 디폴트 스트림은 블로킹 성질을 가지므로 각각의 커널 구동은 다음 커널을 구동하기 전에 완료되어야 한다고 주장할 수 있을 것입니다. 이 역시 맞는 말입니다.\n",
    "\n",
    "[`01-print-numbers`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/05-stream-intro/01-print-numbers.cu)를 수정해서 커널 구동이 비디폴트 스트림에서 일어나게 만드세요. 스트림을 다 사용한 후 필요 없게 되었을 때에는 반드시 제거해야 합니다. 수정된 코드를 아래 셀에서 컴파일하고 실행하세요. 여전히 0부터 4가 출력될 것이지만 반드시 오름차순으로 나타나지는 않을 수도 있습니다. 어떻게 해야 할 지 모르는 경우에는 [솔루션](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/05-stream-intro/solutions/01-print-numbers-solution.cu)을 참고하세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "jixuPwJXZkgw"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o print-numbers-in-streams 05-stream-intro/01-print-numbers.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "8OJbaXQKZkg0"
   },
   "source": [
    "다섯 개의 비디폴트 스트림을 사용하여 다섯 개의 커널을 구동하게 되었는데, 이 커널들의 구동은 순차적일까요 동시적일까요? 스트림에 대하여 여러분이 알게 된 것에 더하여 `printNumber` 커널이 매우 간단한 코드임을 고려하세요. 즉, 병렬로 수행된다고 할 때 커널들이 완료되는 시점이 완전히 일치할까요? \n",
    "\n",
    "가설을 세운 후, 수정된 프로그램의 실행 파일을 위한 nvpp 세션을 열고, 타임라인을 최대화하고 커널 실행을 확대하여 실제 동작을 확인하세요.\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "giUptO-mGIct"
   },
   "source": [
    "### 데이터를 초기화하는 동시성 커널 만들기에 스트림 이용하기\n",
    "\n",
    "지금 사용하고 있는 벡터합 애플리케이션인 [`01-prefetch-check-solution.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/04-prefetch-check/solutions/01-prefetch-check-solution.cu)은 초기화 커널을 세 번 구동합니다.  `vectorAdd` 커널에서 사용되는 세 개의 벡터를 초기화 하기 위해 한 벡터가 하나의 초기화 커널을 갖습니다. 이제 코드를 수정하여 각각의 초기화 커널이 서로 다른 비디폴트 스트림에서 실행되게 하세요. 아래 셀에서 컴파일하고 실행하여 성공 메시지가 나오는 것을 확인하세요. 어떻게 해야 할 지 모르는 경우에는 [솔루션](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/06-stream-init/solutions/01-stream-init-solution.cu)을 참고하세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "LpO1ytjXZkg2"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o init-in-streams 04-prefetch-check/solutions/01-prefetch-check-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "7SOuGaakZkg4"
   },
   "source": [
    "수정된 프로그램의 실행 파일을 위한 nvpp 세션을 열고, 타임라인을 최대화한 후 각 초기화 커널이 고유의 비디폴트 스트림에서 어느 정도 동시에 오버랩되어 실행되는 것을 확인하세요.\n",
    "\n",
    "---\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "Gmxr9Dt8GLPN"
   },
   "source": [
    "## 요약\n",
    "\n",
    "이제 여러분은 아래와 같은 학습 목표를 달성하셨습니다.\n",
    "\n",
    "* **NVIDIA 비주얼 프로파일러(nvvp)**를 이용하여 가속화 애플리케이션의 타임라인을 시각적으로 프로파일링하기\n",
    "* nvpp를 사용하여 가속화 에플리케이션의 최적화 기회를 인식하고 활용하기\n",
    "* CUDA 스트림을 이용하여 가속화 애플리케이션에서 커널을 동시에 실행하기\n",
    "\n",
    "이제 여러분은 CPU-only 애플리케이션을 가속화하고 최적화하기 위해 필요한, 풍부한 기초 도구와 기법을 보유하게 되었습니다. 마지막 실습에서 여러분은 여러분이 배운 모든 것을 [n-body](https://en.wikipedia.org/wiki/N-body_problem) 시뮬레이터의 가속화에 적용할 기회를 갖게 되었습니다. n-body 시뮬레이터는 일군의 물체가 서로 중력을 미치는 상황에서 각 물체의 운동을 예측하는 데에 사용됩니다. \n",
    "\n",
    "---\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "qmTYxTIxGNc9"
   },
   "source": [
    "## 마지막 실습: N-body 시뮬레이터 가속화와 최적화\n",
    "\n",
    "[n-body](https://en.wikipedia.org/wiki/N-body_problem) 시뮬레이터는 일군의 물체가 서로 중력을 미치는 상황에서 각 물체의 운동을 예측합니다.[`01-nbody.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/09-nbody/01-nbody.cu)는 간단하지만 잘 작동하는 3 차원 n-body 시뮬레이터입니다. 커맨드라인 인수를 이용하여 애플리케이션에 서로 영향을 미치는 물체의 갯수를 전달할 수 있습니다. \n",
    "\n",
    "현재는 CPU-ony 형식이고 4096 개의 물체에 대하여 초당 3000 만 번의 상호작용을 계산할 수 있습니다. 여러분이 할 일은 아래와 같습니다.\n",
    "\n",
    "* 시뮬레이션의 정확성을 유지하면서 GPU 가속화하세요.\n",
    "* 반복 작업을 통해 4096(=2<<11) 개의 물체에 대해 300 억 번의 상호작용을 계산하도록 최적화하세요.\n",
    "* 반복 작업을 통해 ~65,000(=2<<15) 개의 물체에 대해 3250 억 번의 상호작용을 계산하도록 최적화하세요.\n",
    "\n",
    "**완료한 다음에는 이 노트북을 여신 브라우저로 돌아가서 Assess 버튼을 누르세요. 시물레이션 정확도를 유지하면서 제시된 목표를 달성했을 경우 여러분은 *CUDA C/C++을 이용한 애플리케이션 가속화 기초*에 대한 역량을 인정하는 인증서를 받으실 수 있습니다.**\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "M0ozqzwZGO-2"
   },
   "source": [
    "### 작업을 위한 고려 사항\n",
    "\n",
    "작업을 수행하기 전에 고려할 사항입니다.\n",
    "\n",
    "* 최초의 수정본에 있어서는 애플리케이션의 로직, 특히 `bodyForce` 함수는 거의 수정이 필요 없습니다. 최대한 쉬운 방법으로 가속화하는 데에 집중하세요.\n",
    "* `randomizeBodies` 함수는 `rand` 함수를 사용하는데 GPU 디바이스는 `rand` 함수를 처리할 수 없으므로 결국 `randomizeBodies`을 가속화할 수 없습니다. 아예 건드리지 마세요.\n",
    "* 코드 베이스는 `main` 안에 for 루프를 가지고 있는데, `bodyForce`에 의해 계산된 물체간의 힘을 물체의 위치와 합치는 역할을 합니다. 합치는 연산은 `bodyForce`가 실행된 후에, 그리고 다음 `bodyForce`를 호출하기 전에 수행되어야 합니다. 어디를 어떻게 병렬화할지 결정할 때 이 점을 유의하세요.\n",
    "* 프로파일에 의한 반복적 접근법을 사용하세요.\n",
    "* 오류 처리 코드를 넣는 것이 요구되지는 않지만 여러분의 코드는 여러분의 책임이니 넣는 것이 도움이 될 것입니다.\n",
    "\n",
    "즐겁게 수행하세요!"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "pg_y2HIOZkg5"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o nbody 09-nbody/01-nbody.cu"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "YtAn4HwUZkg8"
   },
   "outputs": [],
   "source": [
    "!./nbody 11 # This argument is passed as `N` in the formula `2<<N`, to determine the number of bodies in the system"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "2BDQvh5JZkhD"
   },
   "outputs": [],
   "source": [
    "!nvprof ./nbody"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "AOExLUz6ZkhT"
   },
   "source": [
    "---\n",
    "\n",
    "## 고급 주제\n",
    "\n",
    "관심 있는 분들을 위해 이번 섹션에서는 수동 메모리 관리, 비디폴트 스트림을 이용한 커널 동작 오버랩 , 메모리 복사에 대해 조금 더 들어가 보겠습니다. \n",
    "\n",
    "각각의 기법에 대해 학습한 후, n-body 시뮬레이션을 더 개선해 보세요.\n",
    "\n",
    "---\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "PsZo4f1GGV4J"
   },
   "source": [
    "\n",
    "## 수동 메모리 할당과 복사\n",
    "\n",
    "`cudaMallocManaged`와 `cudaMemPrefetchAsync`은 성능이 우수하며 메모리 마이그레이션을 매우 쉽게 만들어 주지만 메모리 할당을 위해 좀 더 수동적인 제어를 할 필요가 있을 수 있습니다. 이는 특히 데이터가 디바이스나 호스트에서만 사용된다는 것을 미리 알 수 있는 경우 더욱 그러한데, 자동적인 요구시 마이그레이션이 필요하지 않을 때 데이터 마이그레이션 비용을 줄일 수 있습니다.\n",
    "\n",
    "또한, 수동 메모리 관리를 이용하면 계산 작업과 데이터 전달을 동시에 수행하기 위해 비디폴트 스트림을 사용할 수 있습니다. 이번 섹션에서 여러분은 기본적인 수동 메모리 할당과 복사 기법을 배우고, 이 기법을 확장하여 계산 작업과 데이터 전달을 동시에 수행하도록 만들어 볼 것입니다.\n",
    "\n",
    "수동 메모리 관리를 위한 CUDA 명령어는 아래와 같습니다.\n",
    "\n",
    "* `cudaMalloc`은 활성화된 GPU에 직접 메모리를 할당합니다. 이는 모든 GPU 페이지 폴트를 제거합니다. 대신에 여기서 리턴된 포인터는 호스트 코드가 사용할 수 없습니다.\n",
    "* `cudaMallocHost`는 CPU에 직접 메모리를 할당합니다. 이는 또한 메모리를 고정(pin)하는데 이는 다른 말로 페이지를 잠근다고도 합니다. 이는 GPU로 또는 GPU로부터의 비동기적 메모리 복사를 가능하게 합니다. 너무 많은 메모리를 고정시키면 CPU 성능을 방해하기 때문에 주의해서 사용해야 합니다. 고정된 메모리는 `cudaFreeHost`로 해제합니다. \n",
    "* `cudaMemcpy`은 메모리를 전달하는 것이 아니고 복사합니다. DtoH와 HtoD 모두 가능합니다.\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "H3CfXv1sGXTp"
   },
   "source": [
    "### 수동 메모리 관리 예제\n",
    "\n",
    "위에서 설명한 CUDA API 호출을 하는 예제 스니펫은 아래와 같습니다.\n",
    "\n",
    "\n",
    "```cpp\n",
    "int *host_a, *device_a;        // Define host-specific and device-specific arrays.\n",
    "cudaMalloc(&device_a, size);   // `device_a` is immediately available on the GPU.\n",
    "cudaMallocHost(&host_a, size); // `host_a` is immediately available on CPU, and is page-locked, or pinned.\n",
    "\n",
    "initializeOnHost(host_a, N);   // No CPU page faulting since memory is already allocated on the host.\n",
    "\n",
    "// `cudaMemcpy` takes the destination, source, size, and a CUDA-provided variable for the direction of the copy.\n",
    "cudaMemcpy(device_a, host_a, size, cudaMemcpyHostToDevice);\n",
    "\n",
    "kernel<<<blocks, threads, 0, someStream>>>(device_a, N);\n",
    "\n",
    "// `cudaMemcpy` can also copy data from device to host.\n",
    "cudaMemcpy(host_a, device_a, size, cudaMemcpyDeviceToHost);\n",
    "\n",
    "verifyOnHost(host_a, N);\n",
    "\n",
    "cudaFree(device_a);\n",
    "cudaFreeHost(host_a);          // Free pinned memory like this.\n",
    "```\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "G5e4CL3PGch5"
   },
   "source": [
    "### 실습: 호스트/디바이스 메모리 수동 할당하기\n",
    "\n",
    "가장 마지막 벡터합 애플리케이션인 [01-stream-init-solution](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/06-stream-init/solutions/01-stream-init-solution.cu)은 `cudaMallocManaged`를 사용하여 메모리를 할당합니다. 할당된 메모리는 우선 커널 초기화를 위한 디바이스에서 사용되고, 다음에는 벡터합을 위한 디바이스에서 사용되며, 마지막으로 호스트로 전달되어 결과를 검증하는 데에 사용합니다. 이는 합당한 접근법이지만 성능에 어떤 영향이 있는지 살펴 보기 위해 수동 메모리 할당과 복사를 실험해 보겠습니다. \n",
    "\n",
    "[`01-stream-init-solution`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/06-stream- init/solutions/01-stream-init-solution.cu)를 수정하여 `cudaMallocManaged`을 사용하지 **않도록** 만드세요. 이를 위해 다음 작업을 수행해야 합니다.\n",
    "\n",
    "* `cudaMallocManaged` 호출을 `cudaMalloc`으로 변경하세요.\n",
    "* 호스트상에서의 검증을 위해 벡터를 하나 더 만드세요. 이것을 만드는 이유는 `cudaMalloc`으로 할당한 메모리가 호스트에게 보이지 않기 때문입니다. 이 호스트 벡터를 `cudaMallocHost`를 이용해서 할당하세요. \n",
    "* `addVectorsInto` 커널 실행이 완료된 후, `cudaMemcpy`를 이용해서 결과 벡터를 `cudaMallocHost`로 만든 호스트상의 벡터로 복사하세요.\n",
    "* `cudaFreeHost`를 이용해서 `cudaMallocHost`로 할당한 메모리를 해제하세요.\n",
    "\n",
    "어떻게 해야 할 지 모르는 경우에는 [솔루션](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/07-manual-malloc/solutions/01-manual-malloc-solution.cu)을 참고하세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "ChBtlJbvZkhU"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o vector-add-manual-alloc 06-stream-init/solutions/01-stream-init-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "aAHeh2RzZkhY"
   },
   "source": [
    "수정이 끝나면 실행 파일을 새로운 nvpp 세션애서 열고 타임라인을 사용하여 아래 작업을 수행하세요.\n",
    "\n",
    "* 타임라인에 더이상 *통합 메모리* 섹션이 없다는 점을 유의하세요.\n",
    "* 이전 수정 버전과 현재 버전의 타임라인을 비교하기 위해 룰러를 사용하여 현재 버전의 `cudaMalloc`과 이전 버전의 `cudaMallocManaged` 간의 실행 시간을 비교하세요.\n",
    "* 현재 버전은 이전 버전보다 초기화 커널 작업이 더 나중에 실행됨에 유의하세요. 타임라인을 실펴 보면 차이는 `cudaMallocHost`의 실행 시간에서 기인함을 알 수 있습니다. 이 점이 바로 메모리 전달과 메모리 복사의 차이점을 극명히 보여 줍니다. 지금 하는 것과 같이 메모리를 복사할 경우 데이터는 서로 다른 두 위치에 존재하게 됩니다. 현재 버전과 같은 호스트 전용 네 번째 벡터의 할당은 이전 버전에서와 같이 세 개의 벡터만을 사용하는 경우보다 약간의 성능 저하를 비용으로 지불해야 합니다.\n",
    "\n",
    "---\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "WCqn-Ii4Gh0S"
   },
   "source": [
    "\n",
    "## 데이터 전달과 코드 실행을 동시에 하기 위해 스트림 사용하기\n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 884,
     "status": "ok",
     "timestamp": 1541022559926,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "DoUlRYGlZkhY",
    "outputId": "e6e8d530-4231-498b-f0e1-d10bf219aaea"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vRhw6_OQ3oY57IIgs8JH2AgZpP9gNJ87ee3IYE3RV8Ox7CLj2RLPAm2YtxyhjfDDBjC6owOOfpKL1JS/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "xu11jOdIZkhc"
   },
   "source": [
    "`cudaMemcpy` 외에 `cudaMemcpyAsync`라는 함수도 있는데 이것은 호스트 메모리가 `cudaMallocHost`에 의해 할당되어 고정된(pinned) 경우에 한하여, 호스트에서 디바이스로, 또는 디바이스에서 호스트로 비동기적 복사를 할 때 사용합니다. \n",
    "\n",
    "커널 실행과 마찬가지로 `cudaMemcpyAsync`도 기본적으로 호스트에 대하여 비동기적입니다. 기본적으로는 디폴트 스트림에서 실행되기 때문에 GPU에서 수행되는 다른 CUDA 연산에 대하여 블로킹 특성을 가지고 있습니다. 하지만 `cudaMemcpyAsync` 함수는 옵션인 다섯 번째 인수로서 비디폴트 스트림을 가질 수 있습니다. 비디폴트 스트림 인수를 설정하면 메모리 전달과 다른 비디폴트 스트림에서 실행되는  CUDA 연산이 동시에 실행되게 만들 수 있습니다. \n",
    "\n",
    "흔히 사용되는 유용한 패턴은 다음 세 가지 동작을 함께 사용하여 메모리 전달과 커널 실행을 동시에 진행되도록 하는 것입니다.\n",
    "\n",
    "* 고정된 호스트 메모리 사용\n",
    "* 비디폴트 스트림에서의 비동기 메모리 복사\n",
    "* 커널을 비디폴트 스트림에서 실행\n",
    "\n",
    "아래 예제 코드에서는 커널의 실행을 위해 메모리 복사가 모두 완료되기를 기다리는 대신에, 필요한 데이터의 일부 세그먼트가 복사되기만 하면 커널 작업을 일부 진행하는 식으로, 일부 복사/일부 실행을 고유의 비디폴트 스트림에서 실행하는 방식으로 되어 있습니다. 이러한 기법을 사용하면 일부 데이터 세그먼트에 대한 작업을 진행되는 것과 동시에 후속 데이터 세그먼트의 복사가 진행되도록 만들 수 있습니다. 이 기법을 사용할 때에는 반복 횟수, 배열의 오프셋 등 세그먼트와 관련된 값이 정확하도록 특히 주의를 기울여야 합니다.\n",
    "\n",
    "\n",
    "```cpp\n",
    "int N = 2<<24;\n",
    "int size = N * sizeof(int);\n",
    "\n",
    "int *host_array;\n",
    "int *device_array;\n",
    "\n",
    "cudaMallocHost(&host_array, size);               // Pinned host memory allocation.\n",
    "cudaMalloc(&device_array, size);                 // Allocation directly on the active GPU device.\n",
    "\n",
    "initializeData(host_array, N);                   // Assume this application needs to initialize on the host.\n",
    "\n",
    "const int numberOfSegments = 4;                  // This example demonstrates slicing the work into 4 segments.\n",
    "int segmentN = N / numberOfSegments;             // A value for a segment's worth of `N` is needed.\n",
    "size_t segmentSize = size / numberOfSegments;    // A value for a segment's worth of `size` is needed.\n",
    "\n",
    "// For each of the 4 segments...\n",
    "for (int i = 0; i < numberOfSegments; ++i)\n",
    "{\n",
    "  // Calculate the index where this particular segment should operate within the larger arrays.\n",
    "  segmentOffset = i * segmentN;\n",
    "\n",
    "  // Create a stream for this segment's worth of copy and work.\n",
    "  cudaStream_t stream;\n",
    "  cudaStreamCreate(&stream);\n",
    "  \n",
    "  // Asynchronously copy segment's worth of pinned host memory to device over non-default stream.\n",
    "  cudaMemcpyAsync(&device_array[segmentOffset],  // Take care to access correct location in array.\n",
    "                  &host_array[segmentOffset],    // Take care to access correct location in array.\n",
    "                  segmentSize,                   // Only copy a segment's worth of memory.\n",
    "                  cudaMemcpyHostToDevice,\n",
    "                  stream);                       // Provide optional argument for non-default stream.\n",
    "                  \n",
    "  // Execute segment's worth of work over same non-default stream as memory copy.\n",
    "  kernel<<<number_of_blocks, threads_per_block, 0, stream>>>(&device_array[segmentOffset], segmentN);\n",
    "  \n",
    "  // `cudaStreamDestroy` will return immediately (is non-blocking), but will not actually destroy stream until\n",
    "  // all stream operations are complete.\n",
    "  cudaStreamDestroy(stream);\n",
    "}\n",
    "```\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "0150s-AVGnmg"
   },
   "source": [
    "### 실습:  커널 실행과 호스트로의 메모리 복사 동시에 실행하기\n",
    "\n",
    "가장 최근 버전의 벡터합 코드인 [`01-manual-malloc-solution.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/07-manual-malloc/solutions/01-manual-malloc-solution.cu)는 모든 벡터합 연산을 GPU에서 수행한 후, 메모리를 호스트로 복사하여 결과를 검증합니다.\n",
    "\n",
    "[`01-manual-malloc-solution.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/07-manual-malloc/solutions/01-manual-malloc-solution.cu)를 수정하되, 네 개의 세그먼트와 비디폴트 스트림을 이용하여 벡터합을 수행해도록 바꾸세요. 벡터합 연산이 끝나기를 기다리지 않고 비동기 메모리 복사가 시작되도록 해야 합니다. 어떻게 해야 할 지 모르는 경우에는 [솔루션](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP-kr/08-overlap-xfer/solutions/01-overlap-xfer-solution.cu)을 참고하세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "-55peqAAZkhe"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o vector-add-manual-alloc 07-manual-malloc/solutions/01-manual-malloc-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "O-Ch-m2yZkhh"
   },
   "source": [
    "수정이 끝나면 실행 파일을 새로운 nvpp 세션애서 열고 타임라인을 사용하여 아래 작업을 수행하세요.\n",
    "\n",
    "* 디바이스에서 호스트로 메모리 전달이 일어나는 시점에 유의하세요. 커널 작업이 완료되기 전인가요 후인가요?\n",
    "* 네 개의 세그먼트의 복사 자체는 서로 오버랩되지 않음에 유의하세요. 서로 다른 비디폴트 스트림에서 실행됨에도 불구하고 주어진 방향(여기에서는 DtoH)으로의 메모리 전달은 하나씩만 일어날 수 있습니다. 여기에서의 성능 이득은 메모리 전달을 일찍 시작한다는 점에 있습니다. 간단한 덧셈보다 훨씬 복잡한 연산을 수행하는 애플리케이션에 있어서는 메모리 복사가 일찍 시작되는 것뿐 아니라 커널 동작과 확실히 오버랩될 것이라는 것을 상상하는 것은 어렵지 않겠지요. "
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "6lOe0bAhGrLJ"
   },
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  }
 ],
 "metadata": {
  "colab": {
   "collapsed_sections": [],
   "name": "03_AC_STREAMS_NVVP-kr.ipynb",
   "provenance": [],
   "version": "0.3.2"
  },
  "kernelspec": {
   "display_name": "Python 2",
   "language": "python",
   "name": "python2"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 2
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython2",
   "version": "2.7.12"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 1
}
